#include<iostream>
#include<stdio.h>
#include<cuda_runtime.h>
#include"freshman.h"

__global__ void matrixdim2_add(float* a,float* b,float* c,int nx,int ny)
{
    int idx=threadIdx.x+blockDim.x*blockIdx.x;
    int idy=threadIdx.y+blockDim.y*blockIdx.y;
    int id=idx+idy*ny;
    if(idx<nx && idy<ny)
    {
        c[id]=a[id]+b[id];
    }
}

void matrixdim2_host(float* a,float* b,float* c,int nx,int ny)
{
    for(int i=0;i<ny;i++)
    {
        for(int j=0;j<nx;j++)
        {
        c[j]=a[j]+b[j];
        }
        c+=nx;
        b+=nx;
        a+=nx;
    }
}

int main()
{
    //初始
    initDevice(0);
    int nx=1<<12;
    int ny=1<<12;
    int nNUM=nx*ny;
    int nByte=sizeof(float)*nNUM;

    //准备host内存
    float* a_h =(float*)malloc(nByte);
    float* b_h =(float*)malloc(nByte);
    float* result_h =(float*)malloc(nByte);
    float* result_GPU_h =(float*)malloc(nByte);
    initialData(a_h,nNUM);
    initialData(b_h,nNUM);
    memset(result_h,0,nByte);
    memset(result_GPU_h,0,nByte);
    matrixdim2_host(a_h,b_h,result_h,nx,ny);

    //准备device内存
    float* a_d = nullptr;
    float* b_d = nullptr;
    float* result_d = nullptr;
    cudaMalloc((float**)&a_d,nByte);
    cudaMalloc((float**)&b_d,nByte);
    cudaMalloc((float**)&result_d,nByte);

    //copy data to device
    cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice);
    cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice);

    //kernel 2d_block and 2d_grid
    double time_begin,time_end;
    time_begin=cpuSecond();
    dim3 block0(32,32);
    dim3 grid0((nx-1)/block0.x+1,(ny-1)/block0.y+1);
    matrixdim2_add<<<grid0,block0>>>(a_d,b_d,result_d,nx,ny);
    time_end=cpuSecond();
    CHECK(cudaDeviceSynchronize());
    printf("2d_block and 2d_grid\n");
    printf("grid_dim(%d,%d,%d)\tblock_dim(%d,%d,%d)\tTime elapsed %f sec\n",grid0.x,grid0.y,grid0.y,block0.x,block0.y,block0.z,time_end-time_begin);
    CHECK(cudaMemcpy(result_GPU_h,result_d,nByte,cudaMemcpyDeviceToHost));
    checkResult(result_h,result_GPU_h,nNUM);

    //kernel 2d_block and 1d_grid
    time_begin=cpuSecond();
    dim3 block1(32,32);
    dim3 grid1((nNUM-1)/(block1.x*block1.y)+1);
    matrixdim2_add<<<grid1,block1>>>(a_d,b_d,result_d,(nNUM-1)/32+1,32);
    time_end=cpuSecond();
    CHECK(cudaDeviceSynchronize());
    printf("2d_block and 1d_grid\n");
    printf("grid_dim(%d,%d,%d)\tblock_dim(%d,%d,%d)\tTime elapsed %f sec\n",grid1.x,grid1.y,grid1.y,block1.x,block1.y,block1.z,time_end-time_begin);
    CHECK(cudaMemcpy(result_GPU_h,result_d,nByte,cudaMemcpyDeviceToHost));
    
    checkResult(result_h,result_GPU_h,nNUM);

    //kernel 1d_block and 2d_grid
    time_begin=cpuSecond();
    dim3 block2(32);
    dim3 grid2((nx-1)/block2.x+1,ny);
    matrixdim2_add<<<grid2,block2>>>(a_d,b_d,result_d,nx,ny);
    time_end=cpuSecond();
    CHECK(cudaDeviceSynchronize());
    printf("1d_block and 2d_grid\n");
    printf("grid_dim(%d,%d,%d)\tblock_dim(%d,%d,%d)\tTime elapsed %f sec\n",grid2.x,grid2.y,grid2.y,block2.x,block2.y,block2.z,time_end-time_begin);
    CHECK(cudaMemcpy(result_GPU_h,result_d,nByte,cudaMemcpyDeviceToHost));
    checkResult(result_h,result_GPU_h,nNUM);

    //kernel 1d_block and 1d_grid
    time_begin=cpuSecond();
    dim3 block3(32);
    dim3 grid3((nNUM-1)/block3.x+1);
    matrixdim2_add<<<grid3,block3>>>(a_d,b_d,result_d,nNUM,1);
    time_end=cpuSecond();
    CHECK(cudaDeviceSynchronize());
    printf("1d_block and 1d_grid\n");
    printf("grid_dim(%d,%d,%d)\tblock_dim(%d,%d,%d)\tTime elapsed %f sec\n",grid3.x,grid3.y,grid3.y,block3.x,block3.y,block3.z,time_end-time_begin);
    CHECK(cudaMemcpy(result_GPU_h,result_d,nByte,cudaMemcpyDeviceToHost));
    checkResult(result_h,result_GPU_h,nNUM);

    //copy data to host
    cudaMemcpy(result_h,result_d,nByte,cudaMemcpyDeviceToHost);
    
    //show result
    showVector(a_h,nNUM);
    showVector(b_h,nNUM);
    showVector(result_h,nNUM);  

    //relase memory
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(result_d);
    cudaDeviceReset();
    free(a_h);
    free(b_h);
    free(result_h);

    return(0);
}
